{
 "cells": [
  {
   "cell_type": "markdown",
   "metadata": {},
   "source": [
    "Before we begin, let us execute the below cell to display information about the NVIDIA® CUDA® driver and the GPUs running on the server by running the `nvidia-smi` command. To do this, execute the cell block below by clicking on it with your mouse, and pressing Ctrl+Enter, or pressing the play button in the toolbar above. You should see some output returned below the grey cell."
   ]
  },
  {
   "cell_type": "code",
   "execution_count": null,
   "metadata": {},
   "outputs": [],
   "source": [
    "!nvidia-smi"
   ]
  },
  {
   "cell_type": "markdown",
   "metadata": {},
   "source": [
    "## Learning objectives\n",
    "The **goal** of this lab is to:\n",
    "\n",
    "- Learn how to run the same code on both a multicore CPU and a GPU using the OpenMP Target programming model\n",
    "- Understand the key directives and steps involved in making a sequential code parallel\n",
    "\n",
    "We do not intend to cover:\n",
    "- Optimization techniques in details\n",
    "\n",
    "\n",
    "# OpenMP Directives\n",
    "- OpenMP was formed in 1997 to focus on vendor-neutral Shared Memory Parallelism.\n",
    "- OpenMP 4.0 in 2013 expanded its focus beyond shared memory parallel computers including accelerators. \n",
    "- The OpenMP 4.0 target construct provides the means to offload data and computation to accelerators.\n",
    "\n",
    "Like OpenACC, OpenMP is directive based. Compiler directives appear as comments in your source code and are ignored by compilers unless you tell them otherwise - usually by specifying the appropriate compiler flag.\n",
    "\n",
    "In this notebook we will  use the OpenMP target construct to offload data and computation to GPU. Multiple compilers are in development to support OpenMP offloading to NVIDIA GPUs. We will be using NVIDIA HPC SDK compiler for this tutorial.\n",
    "\n",
    "\n",
    "## OpenMP Syntax\n",
    "\n",
    "\n",
    "<details>\n",
    "    <summary markdown=\"span\"><b>C/C++ syntax</b></summary>\n",
    "    \n",
    "```#pragma omp directive ``` \n",
    "</details>\n",
    "<br/>\n",
    "    \n",
    "<details>\n",
    "    <summary markdown=\"span\"><b>Fortran syntax</b></summary>\n",
    "    \n",
    "```!$omp directive ``` \n",
    "</details>\n",
    "<br/>\n",
    "\n",
    "**#pragma** in C/C++ and **!$** in Fortran are what's known as a \"compiler hint.\" These are very similar to programmer comments, however, the compiler will  read our pragmas. Pragmas are a way for the programmer to \"guide\" the compiler, without running the chance of damaging the code. If the compiler does not understand the pragma, it can ignore it, rather than throw a syntax error.\n",
    "\n",
    "**omp** is an addition to our pragma,  known as the “sentinel”. It specifies that this is an OpenMP pragma. Any non-OpenMP compiler will ignore this pragma. \n",
    "\n",
    "**directives** are commands in OpenMP that will tell the compiler to do some action. For now, we will only use directives that allow the compiler to parallelize our code.\n",
    "\n",
    "    \n",
    "For beginners new to OpenMP directives, we will introduce some terminologies and concepts before adding ```target``` directives to our code to offload onto GPU computation and data.\n",
    "\n",
    "## OpenMP Fork-Join Model\n",
    "\n",
    "OpenMP uses the fork-join model of parallel execution. All OpenMP programs begin as a single process: the master thread. The master thread executes sequentially until the first parallel region construct is encountered.\n",
    "\n",
    "**FORK**: The master thread then creates a team of parallel threads.The statements in the program that are enclosed by the parallel region construct are then executed in parallel among the various team threads.\n",
    "\n",
    "**JOIN**: When the team threads complete the statements in the parallel region construct, they synchronize and terminate, leaving only the master thread.\n",
    "\n",
    "<img src=\"../../_common/images/openmp_fork_join.png\" width=\"50%\" height=\"50%\">\n",
    "\n",
    "## OpenMP Parallel Region\n",
    "\n",
    "A parallel region is a block of code  executed by multiple threads. This is the fundamental OpenMP parallel construct. When a thread reaches a PARALLEL directive, it creates a team of threads and becomes the master of the team. The master is a member of that team. Starting from the beginning of this parallel region, the code is duplicated, and all threads will execute that code redundantly.There is an implied barrier at the end of a parallel region. Only the master thread continues execution past this point.\n",
    "\n",
    "\n",
    "\n",
    "<details>\n",
    "    <summary markdown=\"span\"><b>C/C++ syntax</b></summary>\n",
    "    \n",
    "    \n",
    "```cpp\n",
    "//Include the header file\n",
    "#include <omp.h>\n",
    "\n",
    " main(int argc, char *argv[]) {\n",
    "\n",
    " int nthreads;\n",
    "\n",
    " /* Fork a team of threads*/\n",
    " #pragma omp parallel\n",
    "   {\n",
    "\n",
    "   /* Obtain and print thread id */\n",
    "   printf(\"Hello World from thread = %d\\n\", omp_get_thread_num());\n",
    "\n",
    "   /* Only master thread does this */\n",
    "   if (omp_get_thread_num() == 0) \n",
    "     {\n",
    "     nthreads = omp_get_num_threads();\n",
    "     printf(\"Number of threads = %d\\n\", nthreads);\n",
    "     }\n",
    "\n",
    "   }  /* All threads join master thread and terminate */\n",
    "\n",
    " }\n",
    "```\n",
    "</details>\n",
    "<br/>\n",
    "\n",
    "\n",
    "<details>\n",
    "    <summary markdown=\"span\"><b>Fortran syntax</b></summary>\n",
    "    \n",
    "```fortran\n",
    "program hello\n",
    "  integer :: omp_rank\n",
    "!$omp parallel private(omp_rank)\n",
    "  omp_rank = omp_get_thread_num()\n",
    "  print *, 'Hello world! by thread ', omp_rank\n",
    "!$omp end parallel\n",
    "end program hello\n",
    " ```\n",
    "</details>\n",
    "<br/>\n",
    "\n",
    "<img src=\"../../_common/images/openmp_parallel_construct.png\" width=\"50%\" height=\"50%\">\n",
    "\n",
    "## OpenMP Data-sharing\n",
    "In OpenMP, several constructs accept clauses that allow the user to control the data sharing. For example, you can use one of the below clauses in a *Parallel* construct.\n",
    "\n",
    "- `private`: Declares variables to be private to each thread in a team. Private copies of the variable are initialized from the original object when entering the region.\n",
    "- `shared`: Shares variables among all the threads in a team.\n",
    "- `default`: Enables you to affect the data-scope attributes of variables.\n",
    "\n",
    "    \n",
    "<details>\n",
    "    <summary markdown=\"span\"><b>C/C++ syntax</b></summary>\n",
    "    \n",
    "```cpp\n",
    "\n",
    "#pragma omp parallel for default(shared) private(dx)\n",
    "{\n",
    "   for (int i=0; i < N; i++){\n",
    "       for (int j=0; j < N; j++){\n",
    "            dx = a[i] + b[j];\n",
    "       }\n",
    "   }  \n",
    "\n",
    "```\n",
    "</details>\n",
    "<br/>\n",
    "\n",
    "\n",
    "<details>\n",
    "    <summary markdown=\"span\"><b>Fortran syntax</b></summary>\n",
    "    \n",
    "\n",
    "```fortran\n",
    "\n",
    "program hello\n",
    "  integer :: omp_rank\n",
    "!$omp parallel default(shared) private(omp_rank)\n",
    "  omp_rank = omp_get_thread_num()\n",
    "  print *, 'Hello world! by thread ', omp_rank\n",
    "!$omp end parallel\n",
    "end program hello\n",
    "\n",
    "```\n",
    "</details>\n",
    "<br/>\n",
    "\n",
    "## OpenMP Work-sharing\n",
    "\n",
    "As described before ```parallel``` construct creates a team of threads, and the execution continues redundantly on all threads of team. Ideally we would need all threads within the team to work share i.e. split the work. A work-sharing construct divides the execution of the enclosed code region among the members of the team that encounter it. Work-sharing constructs do not launch new threads but divide (“workshares”) the iterations of the  loop across the threads in the team . There is no implied barrier upon entry to a work-sharing construct, however there is an implied barrier at the end of a work-sharing construct. \n",
    "\n",
    "There are multiple ways to allow worksharing, the code below makes use of ```for``` to divide the iteration of a loop among threads.\n",
    "\n",
    "\n",
    "<details>\n",
    "    <summary markdown=\"span\"><b>C/C++ syntax</b></summary>\n",
    "    \n",
    "```cpp\n",
    "\n",
    "//Create a team of threads\n",
    "#pragma omp parallel \n",
    "{\n",
    "//workshare this loop across those threads.\n",
    "   #pragma omp for\n",
    "   for (i=0; i < N; i++)\n",
    "     c[i] = a[i] + b[i];\n",
    "\n",
    "   }   /* end of parallel region */\n",
    "\n",
    "```\n",
    "</details>\n",
    "<br/>\n",
    "\n",
    "\n",
    "<details>\n",
    "    <summary markdown=\"span\"><b>Fortran syntax</b></summary>\n",
    "    \n",
    "\n",
    "\n",
    "```fortran\n",
    "\n",
    "!Create a team of threads\n",
    "!$omp parallel\n",
    "!workshare this loop across those threads.\n",
    "   !$omp for\n",
    "    do i=1,N\n",
    "        < loop code >\n",
    "    end do\n",
    "!$omp end parallel\n",
    "\n",
    "```\n",
    "</details>\n",
    "<br/>\n",
    "    \n",
    "<img src=\"../../_common/images/openmp_parallelfor_construct.png\" width=\"50%\" height=\"50%\">\n",
    "\n",
    "\n",
    "\n",
    "## OpenMP Target Offloading\n",
    "\n",
    "By now, you should have got familiar with the OpenMP programming model. Now let us  introduce key directives and construct used to add GPU offloading. \n",
    "\n",
    "\n",
    "### ```target ```\n",
    "\n",
    "```target``` construct consists of a target directive and an execution region. ```target``` directive defines a target region,  a block of computation that operates within a distinct data environment and is intended to be offloaded onto a parallel computation device during execution ( GPU in our case). Data used within the region may be implicitly or explicitly mapped to the device.  OpenMP is allowed within target regions, but only a subset will run well on GPUs.\n",
    "\n",
    "The example below shows the usage of target directive with implicitly mapped data\n",
    "\n",
    "<details>\n",
    "    <summary markdown=\"span\"><b>C/C++ syntax</b></summary>\n",
    "    \n",
    "```cpp\n",
    "while (iter < iter_max )\n",
    "{\n",
    "    error = 0.0;\n",
    "    //Moves this region of code to the GPU and implicitly maps data.\n",
    "    #pragma omp target\n",
    "    {\n",
    "        #pragma omp parallel for reduction(max:error)\n",
    "        for( int j = 1; j < n-1; j++) {\n",
    "            ANew[j] = A [j-1] + A[j+1];\n",
    "        }\n",
    "    }\n",
    "    iter++;\n",
    "}\n",
    "```\n",
    "</details>\n",
    "<br/>\n",
    "\n",
    "\n",
    "<details>\n",
    "    <summary markdown=\"span\"><b>Fortran syntax</b></summary>\n",
    "    \n",
    "```fortran\n",
    "    !Moves this region of code to the GPU and implicitly maps data.\n",
    "    !$omp target\n",
    "        !$omp parallel do\n",
    "        do j=2,N-1 \n",
    "            ANew(j) = A (j-1) + A(j+1)\n",
    "        end do\n",
    "    !$omp end target    \n",
    "```\n",
    "</details>\n",
    "<br/>\n",
    "    \n",
    "###  ```target data``` to explicitly map the data\n",
    "\n",
    "Map a variable to/from the device.Map directive helps developers to explicitly define and reduce data copies. The ```target data``` construct is used to mark such regions.\n",
    "\n",
    "\n",
    " <details>\n",
    " <summary markdown=\"span\"><b>C/C++ syntax</b></summary>\n",
    "\n",
    "```cpp\n",
    "#pragma omp target map(map-type: list)\n",
    "```\n",
    "</details>\n",
    "<br/>\n",
    "\n",
    "\n",
    "<details>\n",
    "    <summary markdown=\"span\"><b>Fortran syntax</b></summary>\n",
    "    \n",
    "```fortran\n",
    "!$omp target data map(map-type: list)\n",
    "```\n",
    "</details>\n",
    "<br/>\n",
    "\n",
    "Examples of mapping data directives are as follows: \n",
    "- `to` (list)\n",
    "    - Allocates memory on the device and copies data in when entering the region, the values are not copied back\n",
    "- `from` (list)\n",
    "    - Allocates memory on the device and copies the data to the host when exiting the region\n",
    "- `alloc` (list)\n",
    "    - Allocates memory on the device. If the data is already present on the device, a reference counter is incremented.\n",
    "\n",
    "\n",
    "<details>\n",
    "    <summary markdown=\"span\"><b>C/C++ syntax</b></summary>\n",
    "    \n",
    "```cpp\n",
    "while (iter < iter_max )\n",
    "{\n",
    "    error = 0.0;\n",
    "    //Moves this region of code to the GPU and explicitly maps data.\n",
    "    #pragma omp target data map(to:A[:n]) map(from:ANew[:n])\n",
    "    {\n",
    "        #pragma omp parallel for reduction(max:error)\n",
    "        for( int j = 1; j < n-1; j++) {\n",
    "            ANew[j] = A [j-1] + A[j+1];\n",
    "        }\n",
    "    }\n",
    "    iter++;\n",
    "}\n",
    "```\n",
    "\n",
    "</details>\n",
    "<br/>\n",
    "\n",
    "\n",
    "<details>\n",
    "    <summary markdown=\"span\"><b>Fortran syntax</b></summary>\n",
    "    \n",
    "```fortran\n",
    "!Moves this region of code to the GPU and explicitly maps data.\n",
    "!$omp target data map(to:A(:)) map(from:ANew(:))\n",
    "        !$omp parallel do\n",
    "        do j=2,N-1 \n",
    "            ANew(j) = A (j-1) + A(j+1)\n",
    "        end do\n",
    "!$omp end target data\n",
    "\n",
    "\n",
    "```\n",
    "</details>\n",
    "<br/>\n",
    "    \n",
    "### ```teams``` directive\n",
    "```teams``` directve creates a league of thread teams where the master thread of each team executes the region. Each of these master threads executes sequentially. In other words, teams directive spawn one or more thread teams with the same number of threads. The execution continues on the master threads of each team (redundantly). There is no synchronization allowed between teams. \n",
    "\n",
    "OpenMP calls that somewhere a team, which might be a thread on the CPU or maying a CUDA threadblock or OpenCL workgroup. It will choose how many teams to create based on where you're running, only a few on a CPU (like 1 per CPU core) or lots on a GPU (1000's possibly). ```teams``` allow OpenMP code to scale from small CPUs to large GPUs because each one works completely independently of the other ```teams```.\n",
    "\n",
    "<img src=\"../../_common/images/openmp_target_teams.png\" width=\"50%\" height=\"50%\">\n",
    "\n",
    "### ```distribute``` \n",
    "There's a good chance that we don't want the loop to be run redundantly in every master thread of ```teams``` though, that seems wasteful and potentially dangerous. With the usage of ```distribute``` construct the iterations of the next loop are broken into groups that are *distributed* to the master threads of the teams. The iterations are distributed statically and there is no guarantee about the order teams will execute. Also it does not generate parallelism/worksharing within the thread teams.\n",
    "\n",
    "<img src=\"../../_common/images/openmp_target_distribute.png\" width=\"50%\" height=\"50%\">\n",
    "\n",
    "The example below of simple stencil code shows the usage of ```distribute``` along with ```team```:\n",
    "\n",
    "    \n",
    "<details>\n",
    "    <summary markdown=\"span\"><b>C/C++ syntax</b></summary>\n",
    "    \n",
    "```cpp\n",
    "#pragma omp target teams distribute \n",
    "    for( int j = 1; j < n-1; j++) {\n",
    "       for( int i = 1; i < m-1; i++) {\n",
    "            Anew[j][i] = 0.25 * ( A[j][i+1] + A [j][i-1]\n",
    "                            + A[j-1][i] + A[j+1][i]);\n",
    "            error = fmax (error, fabs(Anew[j][i] - A[j][i]));\n",
    "    }}\n",
    "```\n",
    "</details>\n",
    "<br/>\n",
    "\n",
    "\n",
    "<details>\n",
    "    <summary markdown=\"span\"><b>Fortran syntax</b></summary>\n",
    "\n",
    "```fortran\n",
    "    !$omp target teams distribute\n",
    "        do j=2,N-1 \n",
    "           do i=2,N-1 \n",
    "             ANew(j,i) = 0.25 * (A(j-1,i)  +  A(j+1,i) + A(j,i-1) + A(j,i+1))\n",
    "           enddo\n",
    "        enddo\n",
    "    !$omp end target    \n",
    "```\n",
    "</details>\n",
    "<br/>\n",
    "\n",
    "    \n",
    "<img src=\"../../_common/images/openmp_teams.png\" width=\"80%\" height=\"80%\">\n",
    "\n",
    "\n",
    "### Work sharing to improve parallelism\n",
    "\n",
    "As shown in the image, only the master thread performs the computation which is not so optimal in the case of GPU architecture. To solve this problem we will make use of work-sharing as we did before. When any team encounters a worksharing construct, the work inside the construct is divided among the members of the team, and executed cooperatively instead of being executed by every thread. There are many work-sharing constructs defined, one is using `teams distribute` construct: \n",
    "\n",
    "    \n",
    "<img src=\"../../_common/images/openmp_teams_for.png\" width=\"80%\" height=\"80%\">\n",
    "    \n",
    "    \n",
    "<details>\n",
    "    <summary markdown=\"span\"><b>C/C++ syntax</b></summary>\n",
    "    \n",
    "```cpp\n",
    "#pragma omp target teams distribute parallel for\n",
    "    for( int j = 1; j < n-1; j++) {\n",
    "       for( int i = 1; i < m-1; i++) {\n",
    "            Anew[j][i] = 0.25 * ( A[j][i+1] + A [j][i-1]\n",
    "                            + A[j-1][i] + A[j+1][i]);\n",
    "            error = fmax (error, fabs(Anew[j][i] - A[j][i]));\n",
    "    }}\n",
    "```\n",
    "</details>\n",
    "<br/>\n",
    "\n",
    "\n",
    "<details>\n",
    "    <summary markdown=\"span\"><b>Fortran syntax</b></summary>\n",
    "\n",
    "```fortran\n",
    "    !$omp target teams distribute parallel do\n",
    "        do j=2,N-1\n",
    "           do i=2,N-1 \n",
    "             ANew(j,i) = 0.25 * (A(j-1,i)  +  A(j+1,i) + A(j,i-1) + A(j,i+1))\n",
    "           enddo\n",
    "        enddo\n",
    "    !$omp end target    \n",
    "```\n",
    "</details>\n",
    "<br/>\n",
    "\n",
    "In the above examples, the programmer explicitly requests the steps the compiler should take to map parallelism to the target architecture. Another way to expose more parallelism in a program is to allow a compiler to do the mapping onto the target architectures. The HPC compilers' implementation of the loop supports this descriptive model. In the below examples, the programmer specifies the loop regions to be parallelized by the compiler, and the compilers parallelize the loop across teams and threads using the `teams loop` construct, which is a shortcut for specifying a teams construct containing a loop construct and no other statements:\n",
    "\n",
    "        \n",
    "<details>\n",
    "    <summary markdown=\"span\"><b>C/C++ syntax</b></summary>\n",
    "    \n",
    "```cpp\n",
    "#pragma omp target teams loop reduction(max:error) \n",
    "for( int j = 1; j < n-1; j++) {\n",
    "  #pragma omp loop reduction(max:error)\n",
    "  for( int i = 1; i < m-1; i++ ) {\n",
    "      Anew[j][i] = 0.25f * ( A[j][i+1] + A[j][i-1]\n",
    "                            + A[j-1][i] + A[j+1][i]);\n",
    "      error = fmaxf( error, fabsf(Anew[j][i]-A[j][i]));\n",
    "  }\n",
    "}\n",
    "```\n",
    "</details>\n",
    "<br/>\n",
    "\n",
    "<details>\n",
    "    <summary markdown=\"span\"><b>Fortran syntax</b></summary>\n",
    "\n",
    "```fortran   \n",
    "!$omp target teams loop \n",
    "do n1loc_blk = 1, n1loc_blksize\n",
    "  do igp = 1, ngpown \n",
    "    do ig_blk = 1, ig_blksize \n",
    "      do ig = ig_blk, ncouls, ig_blksize\n",
    "        do n1_loc = n1loc_blk, ntband_dist, n1loc_blksize\n",
    "          !expensive computation codes           \n",
    "        enddo \n",
    "      enddo \n",
    "    enddo \n",
    "  enddo \n",
    "enddo\n",
    "```\n",
    "</details>\n",
    "<br/>\n",
    "\n",
    "Moreover, further tuning when using a `loop` construct can be done with the `bind` clause, where binding can be one of `teams`, `parallel`, or `thread`. For more information, please visit (OpenMP documentation)[https://www.openmp.org/spec-html/5.1/openmpsu51.html].\n",
    "\n",
    "## Atomic Construct\n",
    "\n",
    "In the code, you will also require one more construct, which will help you get the right results. OpenMP atomic construct ensures that a particular variable is accessed and/or updated atomically to prevent indeterminate results and race conditions. In other words, it prevents one thread from stepping on the toes of other threads due to accessing a variable simultaneously, resulting in different results run-to-run. For example, if we want to count the number of elements that have a value greater than zero, we could write the following:\n",
    "\n",
    "\n",
    "<details>\n",
    "    <summary markdown=\"span\"><b>C/C++ syntax</b></summary>\n",
    "\n",
    "\n",
    "```cpp\n",
    "if ( val > 0 )\n",
    "{\n",
    "  #pragma omp atomic\n",
    "  {\n",
    "    cnt++;\n",
    "  }\n",
    "}\n",
    "```\n",
    "</details>\n",
    "<br/>\n",
    "\n",
    "\n",
    "<details>\n",
    "    <summary markdown=\"span\"><b>Fortran syntax</b></summary>\n",
    "    \n",
    "```fortran\n",
    "if(r<cut)then\n",
    "  !$omp atomic\n",
    "    cnt = cnt + 1\n",
    "```\n",
    "\n",
    "</details>\n",
    "<br/>"
   ]
  },
  {
   "cell_type": "markdown",
   "metadata": {},
   "source": [
    "Now, lets start modifying the original code and add the OpenMP directives. Click on the <b>[C/C++ version](../source_code/rdf.cpp)</b> or the <b>[Fortran version](../source_code/rdf.f90)</b> links, and <mark>start modifying the C or Fortran version of the RDF code.  Without changing the orginal code, you will not get the expected outcome after running the below cells. </mark> Remember to **SAVE** your code after changes, before running the below cells."
   ]
  },
  {
   "cell_type": "markdown",
   "metadata": {},
   "source": [
    "### Compile and Run for Multicore\n",
    "\n",
    "Having added OpenMP directives, let us compile the code. We will be using NVIDIA HPC SDK compiler for this exercise. The flags used for enabling OpenMP target offloading are as follows:\n",
    "\n",
    "<!--\n",
    "**-fopenmp** : This flag will give tell compiler to parse and act on OpenMP directive.\n",
    "**-fopenmp-targets** : This flag allows us to compile our code for a specific target parallel hardware. Without this flag, the code will be compiled for multicore execution.\n",
    "-->\n",
    "\n",
    "`-mp=gpu|multicore` : Select the target device for all parallel programming paradigms used (OpenACC, OpenMP, Standard Languages)\n",
    "- `gpu`             Globally set the target device to an NVIDIA GPU\n",
    "- `multicore`       Globally set the target device to the host CPU\n",
    "\n",
    "**NOTE:** `-Minfo=mp` enables OpenMP information.\n",
    "\n",
    "After running the cells, make sure to check the output first. You can inspect part of the compiler feedback for C or Fortran version and see what it's telling us (your compiler feedback will be similar to the below)."
   ]
  },
  {
   "cell_type": "markdown",
   "metadata": {
    "jp-MarkdownHeadingCollapsed": true
   },
   "source": [
    "### <mark>Compile the code for multicore (C/C++)</mark>"
   ]
  },
  {
   "cell_type": "code",
   "execution_count": null,
   "metadata": {},
   "outputs": [],
   "source": [
    "#Compile the code for multicore (C/C++)\n",
    "!cd ../source_code && echo \"compiling C/C++ version .. \" && nvc++ -mp=multicore -Minfo=mp -I/opt/nvidia/hpc_sdk/Linux_x86_64/23.5/cuda/11.8/include -o rdf_c rdf.cpp && ./rdf_c && cat Pair_entropy.dat"
   ]
  },
  {
   "cell_type": "markdown",
   "metadata": {},
   "source": [
    "Compiler Feedback (C/C++ version):\n",
    "\n",
    "You can see from *Line 174* that it is generating a multicore code `174, Generating Multicore code`.\n",
    "\n",
    "<img src=\"../../_common/images/openmp_feedback_multicore.png\">\n",
    "    \n",
    "The output should be the following:\n",
    "\n",
    "```\n",
    "s2 value is -2.43191\n",
    "s2bond value is -3.87014\n",
    "```\n",
    "Now, let's profile the code."
   ]
  },
  {
   "cell_type": "code",
   "execution_count": null,
   "metadata": {},
   "outputs": [],
   "source": [
    "#profile and see output of nvptx (C/C++ version)\n",
    "!cd ../source_code && nsys profile -t nvtx --stats=true --force-overwrite true -o rdf_multicore_c ./rdf_c"
   ]
  },
  {
   "cell_type": "markdown",
   "metadata": {},
   "source": [
    "Let's checkout the profiler's report.  Download and save the report file by holding down <mark>Shift</mark> and <mark>right-clicking</mark> the [C/C++ version](../source_code/rdf_multicore_c.nsys-rep) then choosing <mark>save Link As</mark> Once done, open it via the GUI. Have a look at the example expected profiler report below:\n",
    "\n",
    "**Example screenshot (C/C++ code)**\n",
    "\n",
    "<img src=\"../../_common/images/openmp_multicore.png\">\n",
    "\n",
    "Feel free to checkout the solutions for [C/C++](../source_code/SOLUTION/rdf_offload.cpp) version to help you understand better."
   ]
  },
  {
   "cell_type": "markdown",
   "metadata": {
    "jp-MarkdownHeadingCollapsed": true
   },
   "source": [
    "### <mark>Compile the code for multicore (Fortran)</mark>"
   ]
  },
  {
   "cell_type": "code",
   "execution_count": null,
   "metadata": {},
   "outputs": [],
   "source": [
    "#Compile the code for multicore (Fortran)\n",
    "!cd ../source_code && echo \"compiling Fortran version .. \" && nvfortran -mp=multicore -Minfo=mp -o rdf_f rdf.f90 -lnvhpcwrapnvtx && ./rdf_f && cat Pair_entropy.dat"
   ]
  },
  {
   "cell_type": "markdown",
   "metadata": {},
   "source": [
    "Compiler Feedback (Fortran version):\n",
    "\n",
    "You can see from *Line 98* that it is generating a multicore code `98, Generating Multicore code`.\n",
    "\n",
    "```\n",
    "\tMulticore output\n",
    "\t\trdf:\n",
    "         98, !$omp target teams distribute parallel do\n",
    "         98, Generating Multicore code\n",
    "         99, Loop parallelized across teams and threads, schedule(static)\n",
    "```\n",
    "The output should be the following:\n",
    "\n",
    "```\n",
    "s2      :    -2.452690945278331     \n",
    "s2bond  :    -24.37502820694527    \n",
    "```\n",
    "Now, let's profile the code."
   ]
  },
  {
   "cell_type": "code",
   "execution_count": null,
   "metadata": {},
   "outputs": [],
   "source": [
    "#profile and see output of nvptx (Fortran version)\n",
    "!cd ../source_code && nsys profile -t nvtx --stats=true --force-overwrite true -o rdf_multicore_f ./rdf_f"
   ]
  },
  {
   "cell_type": "markdown",
   "metadata": {},
   "source": [
    "Let's checkout the profiler's report.  Download and save the report file by holding down <mark>Shift</mark> and <mark>right-clicking</mark> the [Fortran version](../source_code/rdf_multicore_f.nsys-rep) then choosing <mark>save Link As</mark> Once done, open it via the GUI. Have a look at the example expected profiler report below:\n",
    "\n",
    "**Example screenshot (Fortran code)**\n",
    "\n",
    "<img src=\"../../_common/images/f_openmp_multicore.png\">\n",
    "\n",
    "\n",
    "Feel free to checkout the solutions for [Fortran](../source_code/SOLUTION/rdf_offload.f90) version to help you understand better."
   ]
  },
  {
   "cell_type": "markdown",
   "metadata": {},
   "source": [
    "### Compile and Run for an NVIDIA GPU\n",
    "\n",
    "Without changing the code now let us try to recompile the code for NVIDIA GPU and rerun.\n",
    "The only difference is that now we pass `gpu` value to the `-mp` compiler option.`-mp=gpu`.  **Understand and analyze** the code present at <b>[C/C++ version](../source_code/SOLUTION/rdf_offload.cpp)</b> and/or the <b>[Fortran version](../source_code/SOLUTION/rdf_offload.f90)</b> .\n",
    "\n",
    "Open the downloaded files for inspection. Once done, modify your code by using `teams loop` construct. Click on the <b>[C/C++ version](../source_code/rdf.cpp)</b> or the <b>[Fortran version](../source_code/rdf.f90)</b> links, and <mark>start modifying the C or Fortran version of the RDF code.  Without changing the orginal code, you will not get the expected outcome after running the below cells. </mark> Remember to **SAVE** your code after changes, before running below cells. \n",
    "\n",
    "After running the cells, make sure to check the output first. You can inspect part of the compiler feedback for C or Fortran version and see what it's telling us (your compiler feedback will be similar to the one below)."
   ]
  },
  {
   "cell_type": "markdown",
   "metadata": {
    "jp-MarkdownHeadingCollapsed": true
   },
   "source": [
    "### <mark>Compile the code for GPU (C/C++)</mark>"
   ]
  },
  {
   "cell_type": "code",
   "execution_count": null,
   "metadata": {},
   "outputs": [],
   "source": [
    "#compile for Tesla GPU (C/C++ version)\n",
    "!cd ../source_code && echo \"compiling C/C++ version .. \" && nvc++ -mp=gpu -Minfo=mp -o rdf_c rdf.cpp && ./rdf_c && cat Pair_entropy.dat"
   ]
  },
  {
   "cell_type": "markdown",
   "metadata": {},
   "source": [
    "Compiler Feedback (C/C++ version):\n",
    "\n",
    "Inspect the compiler feedback (you should get a similar output as below) and you can see below: \n",
    "\n",
    "- *Line 86* shows variables mapped to the device\n",
    "- *Line 172* shows the GPU kernel is generated `Generating \"nvkernel__Z8pair_gpuPKdS0_S0_Pjiidddi_F1L174_1\" GPU kernel`\n",
    "\n",
    "<!--<img src=\"../../_common/images/openmp_feedback.png\">-->\n",
    "```\n",
    "main:\n",
    "     86, Generating map(tofrom:h_g2[:nbin],h_x[:numatm*nconf],h_y[:numatm*nconf],h_z[:numatm*nconf]) \n",
    "pair_gpu(const double *, const double *, const double *, unsigned int *, int, int, double, double, double, int):\n",
    "    172, #omp target teams loop\n",
    "        172, Generating \"nvkernel__Z8pair_gpuPKdS0_S0_Pjiidddi_F1L172_2\" GPU kernel\n",
    "             Generating NVIDIA GPU code\n",
    "          172, Loop parallelized across teams /* blockIdx.x */\n",
    "          175, Loop parallelized across threads(128) /* threadIdx.x */\n",
    "        172, Generating Multicore code\n",
    "          172, Loop parallelized across threads\n",
    "    172, Generating implicit map(tofrom:d_g2[:]) \n",
    "         Generating implicit allocate(d_x[:]) [if not already present]\n",
    "         Generating implicit copyin(d_x[numatm*frame:*]) [if not already present]\n",
    "         Generating implicit allocate(d_y[:],d_z[:]) [if not already present]\n",
    "         Generating implicit copyin(d_y[numatm*frame:*],d_z[numatm*frame:*]) [if not already present]\n",
    "    175, Loop is parallelizable\n",
    "```\n",
    "    \n",
    "The output should be the following:\n",
    "\n",
    "```\n",
    "s2 value is -2.43191\n",
    "s2bond value is -3.87014\n",
    "```\n",
    "It is very important to inspect the feedback to ensure the compiler is doing what you have asked. Now, let's profile the code."
   ]
  },
  {
   "cell_type": "code",
   "execution_count": null,
   "metadata": {},
   "outputs": [],
   "source": [
    "#profile and see output of nvptx (C/C++)\n",
    "!cd ../source_code && nsys profile -t nvtx,cuda --stats=true --force-overwrite true -o rdf_gpu_c ./rdf_c"
   ]
  },
  {
   "cell_type": "markdown",
   "metadata": {},
   "source": [
    "Download and save the report file by holding down <mark>Shift</mark> and <mark>right-clicking</mark> the [C/C++ version](../source_code/rdf_gpu_c.nsys-rep) then choosing <mark>save Link As</mark> Once done, open it via the GUI. Have a look at the example expected profiler report below:\n",
    "\n",
    "**Example screenshot (C/C++ code)**\n",
    "\n",
    "<img src=\"../../_common/images/openmp_gpu.png\">\n",
    "\n",
    "\n",
    "If you expand the CUDA row (_Timeline View_), you can see memory movements as well as kernels. Checkout the NVTX row and compare the execution time for the `Pair_Calculation` for the multicore version and the GPU offload version. In the *example screenshot*, we were able to reduce the timing significantly.\n",
    "\n",
    "Feel free to checkout the solutions for [C/C++](../source_code/SOLUTION/rdf_offload_loop.cpp) version to help you understand better.\n"
   ]
  },
  {
   "cell_type": "markdown",
   "metadata": {
    "jp-MarkdownHeadingCollapsed": true
   },
   "source": [
    "### <mark>Compile the code for GPU (Fortran)</mark>"
   ]
  },
  {
   "cell_type": "code",
   "execution_count": null,
   "metadata": {},
   "outputs": [],
   "source": [
    "#compile for Tesla GPU (Fortran version)\n",
    "!cd ../source_code && echo \"compiling Fortran version .. \" && nvfortran -mp=gpu -Minfo=mp -o rdf_f rdf.f90 -lnvhpcwrapnvtx && ./rdf_f && cat Pair_entropy.dat"
   ]
  },
  {
   "cell_type": "markdown",
   "metadata": {},
   "source": [
    "Compiler Feedback (Fortran version):\n",
    "\n",
    "Inspect the compiler feedback (you should get a similar output as below) and you can see below: \n",
    "\n",
    "- *Line 94* shows variables mapped to the device\n",
    "- *Line 98* shows the GPU kernel is generated\n",
    "\n",
    "```\n",
    "rdf:\n",
    "     98, !$omp target teams loop\n",
    "         94, Generating map(tofrom:g(:),x(x$sd9:(x$sd9-1)+x$sd9,x$sd9:(x$sd9-1)+x$sd9),\n",
    "y(y$sd8:(y$sd8-1)+y$sd8,y$sd8:(y$sd8-1)+y$sd8),z(z$sd7:(z$sd7-1)+z$sd7,z$sd7:(z$sd7-1)+z$sd7)) \n",
    "         98, Generating \"nvkernel_MAIN__F1L98_2\" GPU kernel\n",
    "             Generating NVIDIA GPU code\n",
    "           99, Loop parallelized across teams ! blockidx%x\n",
    "          101, Loop parallelized across threads(128) ! threadidx%x\n",
    "         98, Generating Multicore code\n",
    "           99, Loop parallelized across threads\n",
    "     98, Generating implicit map(tofrom:g(:),x(:,:),z(:,:),y(:,:)) \n",
    "    101, Loop is parallelizable\n",
    "\n",
    "```\n",
    "    \n",
    "The output should be the following:\n",
    "\n",
    "```\n",
    "s2      :    -2.452690945278331     \n",
    "s2bond  :    -24.37502820694527    \n",
    "```\n",
    "It is very important to inspect the feedback to ensure the compiler is doing what you have asked. Now, let's profile the code."
   ]
  },
  {
   "cell_type": "code",
   "execution_count": null,
   "metadata": {},
   "outputs": [],
   "source": [
    "#profile and see output of nvptx (Fortran)\n",
    "!cd ../source_code && nsys profile -t nvtx,cuda --stats=true --force-overwrite true -o rdf_gpu_f ./rdf_f"
   ]
  },
  {
   "cell_type": "markdown",
   "metadata": {},
   "source": [
    "Download and save the report file by holding down <mark>Shift</mark> and <mark>right-clicking</mark> the  [Fortran version](../source_code/rdf_gpu_f.nsys-rep) then choosing <mark>save Link As</mark> Once done, open it via the GUI. Have a look at the example expected profiler report below:\n",
    "\n",
    "**Example screenshot (Fortran code)**\n",
    "\n",
    "<img src=\"../../_common/images/f_openmp_gpu.png\">\n",
    "\n",
    "\n",
    "If you expand the CUDA row (_Timeline View_), you can see memory movements as well as kernels. Checkout the NVTX row and compare the execution time for the `Pair_Calculation` for the multicore version and the GPU offload version. In the *example screenshot*, we were able to reduce the timing significantly.\n",
    "\n",
    "Feel free to checkout the solutions for [Fortran](../source_code/SOLUTION/rdf_offload_loop.f90) version to help you understand better."
   ]
  },
  {
   "cell_type": "markdown",
   "metadata": {},
   "source": [
    "# OpenMP Analysis\n",
    "\n",
    "**Usage Scenarios**\n",
    "- Legacy codes with a sizeable codebase need to be ported to GPUs with minimal code changes to sequential code. If a compiler doesn’t understand your directives, the code will still compile sequentially. So you have the benefit of maintaining a single code base.\n",
    "- Developers want to see if the code structure favours  GPU SIMD/SIMT style or as we say test the waters before moving a large piece of code to a GPU.\n",
    "\n",
    "\n",
    "**Limitations/Constraints**\n",
    "- Directive-based programming model like OpenMP depends on a compiler to understand and convert your sequential code to CUDA constructs. OpenMP compiler with target offload support is evolving and they it cannot match the best performance that says using CUDA C constructs directly can give. Things like controlling execution at warp level or limiting the register counts etc are some examples.\n",
    "    \n",
    "**Which Compilers Support OpenMP on GPU?**\n",
    "As of March 2020 here are the compilers that support OpenMP on GPU:\n",
    "\n",
    "| Compiler | Latest Version | Maintained by | Full or Partial Support |\n",
    "| --- | --- | --- | --- |\n",
    "| GCC | 12 | Mentor Graphics | 4.5 and 5.0 partial spec supported |\n",
    "| CCE| latest | Cray | 4.5 partial spec supported | \n",
    "| XL | latest | IBM | 4.5 partial spec supported |\n",
    "| Clang | 13.0 | Community | 4.5 partial spec supported |\n",
    "| HPC SDK | 22.11 | NVIDIA HPC SDK | 5.0 spec supported |"
   ]
  },
  {
   "cell_type": "markdown",
   "metadata": {},
   "source": [
    "## Post-Lab Summary\n",
    "\n",
    "If you would like to download this lab for later viewing, it is recommended you go to your browser's file menu (not the Jupyter notebook file menu) and save the complete web page.  This will ensure the images are copied down as well. You can also execute the following cell block to create a zip file of the files you have been working on, and download it with the link below."
   ]
  },
  {
   "cell_type": "code",
   "execution_count": null,
   "metadata": {},
   "outputs": [],
   "source": [
    "%%bash\n",
    "cd ..\n",
    "rm -f _files.zip\n",
    "zip -r _files.zip *"
   ]
  },
  {
   "cell_type": "markdown",
   "metadata": {},
   "source": [
    "**After** executing the above zip command, you should be able to download and save the zip file by holding down <mark>Shift</mark> and <mark>right-clicking</mark> [Here](../_files.zip) then choosing <mark>save Link As</mark>.\n",
    "\n",
    "<!--\n",
    "**IMPORTANT**: If you would like to continue and optimize this application further with OpenMP, please click on the **NEXT** button, otherwise click on **HOME** to go back to the main notebook for *N ways of GPU programming for MD* code.\n",
    "\n",
    "\n",
    "**IMPORTANT**: Please click on **HOME** to go back to the main notebook for *N ways of GPU programming for MD* code.\n",
    "\n",
    "-----\n",
    "\n",
    "# <p style=\"text-align:center;border:3px; border-style:solid; border-color:#FF0000  ; padding: 1em\"> <a href=../../../nways_MD_start.ipynb>HOME</a></p>\n",
    "\n",
    "-----\n",
    "\n",
    "-->\n",
    "\n",
    "# Links and Resources\n",
    "[OpenMP Programming Model](https://computing.llnl.gov/tutorials/openMP/)\n",
    "\n",
    "[OpenMP Target Directive](https://www.openmp.org/wp-content/uploads/openmp-examples-4.5.0.pdf)\n",
    "\n",
    "[NVIDIA Nsight System](https://docs.nvidia.com/nsight-systems/)\n",
    "\n",
    "\n",
    "**NOTE**: To be able to see the Nsight Systems profiler output, please download the latest version of Nsight Systems from [here](https://developer.nvidia.com/nsight-systems).\n",
    "\n",
    "Don't forget to check out additional [Open Hackathons Resources](https://www.openhackathons.org/s/technical-resources) and join our [OpenACC and Hackathons Slack Channel](https://www.openacc.org/community#slack) to share your experience and get more help from the community.\n",
    "\n",
    "--- "
   ]
  },
  {
   "cell_type": "markdown",
   "metadata": {},
   "source": [
    "## Licensing \n",
    "\n",
    "Copyright © 2022 OpenACC-Standard.org.  This material is released by OpenACC-Standard.org, in collaboration with NVIDIA Corporation, under the Creative Commons Attribution 4.0 International (CC BY 4.0). These materials may include references to hardware and software developed by other entities; all applicable licensing and copyrights apply."
   ]
  }
 ],
 "metadata": {
  "kernelspec": {
   "display_name": "Python 3 (ipykernel)",
   "language": "python",
   "name": "python3"
  },
  "language_info": {
   "codemirror_mode": {
    "name": "ipython",
    "version": 3
   },
   "file_extension": ".py",
   "mimetype": "text/x-python",
   "name": "python",
   "nbconvert_exporter": "python",
   "pygments_lexer": "ipython3",
   "version": "3.8.10"
  }
 },
 "nbformat": 4,
 "nbformat_minor": 4
}
